
#define CurrentID ((get_group_id(1)*get_num_groups(0)*get_num_groups(2)) + (get_group_id(2)*get_num_groups(0)) + get_group_id(0))
#define BlockNum 16


__kernel void f_ParticleDomain(__global float4* particle, __global uchar* Voxel, uint eNum, float Raduis){
	const size_t id = CurrentID;
	if(id<eNum){
		float3 p = particle[id].xyz;
		if(fabs(p.x) < 32767 && fabs(p.y) < 32767 && fabs(p.z) < 32767 ){
			//ushort3 offset = (ushort3)(32767, 32767, 32767);
			ushort3 tmmpHash =  (ushort3)(p.x+32767, p.y+32767, p.z+32767);

			tmmpHash.x -= get_group_id(0)/3;
			tmmpHash.y -= get_group_id(1)/3;
			tmmpHash.z -= get_group_id(2)/3;

			uchar offset = tmmpHash.z - (tmmpHash.z/8)*8;
			if(offset>=8) printf("err offset %u", offset);
			size_t index = (8191*65536*tmmpHash.y) + (8191*tmmpHash.z) + tmmpHash.z/8;
			Voxel[index] |= 1<<offset;
		}
	}

}

